#ifndef	tensor_included
#define	tensor_included

template<class	T,	uint32_t	N>
struct	tensor {
    T	*p;
    tensor() {
        cudaMallocManaged(&p,	N*sizeof(T));
        zero();
    }
    ~tensor() {
        cudaFree(p);
    }
    void	save(FILE	*F) {
        fwrite(p,N*sizeof(T),1,F);
    }
    bool	load(FILE	*F) {
        return fread(p,N*sizeof(T),1,F)==1;
    }
    uint32_t	size(void) {
        return	N;
    }
    void	zero(void) {
        cudaMemset(p,0,N*sizeof(T));
    }
};

__global__	void	_scalehalf(__half	*inp,	half8	scale) {
    uint32_t	id=(blockIdx.x*blockDim.x+threadIdx.x)<<3;
    *(half8*)(inp+id)=_hmul8(scale,*(half8*)(inp+id));
}

template<uint32_t	N>
static	inline	void	scalehalf(tensor<__half,N>	&inp,	float	s) {
    const	uint32_t	warp=warp_size(8,N);
    half8	t;
    t.x.x=t.x.y=t.y.x=t.y.y=t.z.x=t.z.y=t.w.x=t.w.y=__float2half_rn(s);
    _scalehalf<<<N/8/warp,warp>>>(inp.p,t);
}

__global__	void	_float2half(float	*inp, __half	*out) {
    uint32_t	id=(blockIdx.x*blockDim.x+threadIdx.x)<<3;
    *(half8*)(out+id)=_f2h8(*(float8*)(inp+id));
}

template<uint32_t	N>
static	inline	void	float2half(tensor<float,N>	&inp,	tensor<__half,N>	&out) {
    const	uint32_t	warp=warp_size(8,N);
    _float2half<<<N/8/warp,warp>>>(inp.p,out.p);
}

__global__	void	_float2char2half(float	*inp, __half	*out) {
    uint32_t	id=(blockIdx.x*blockDim.x+threadIdx.x)<<3;
    *(half8*)(out+id)=_f2h8(_c2f8(_f2c8(*(float8*)(inp+id))));
}

template<uint32_t	N>
static	inline	void	float2char2half(tensor<float,N>	&inp,	tensor<__half,N>	&out) {
    const	uint32_t	warp=warp_size(8,N);
    _float2char2half<<<N/8/warp,warp>>>(inp.p,out.p);
}

__global__	void	_half2float(__half	*inp, float	*out) {
    uint32_t	id=(blockIdx.x*blockDim.x+threadIdx.x)<<3;
    *(float8*)(out+id)=_h2f8(*(half8*)(inp+id));
}

template<uint32_t	N>
static	inline	void	half2float(tensor<__half,N>	&inp,	tensor<float,N>	&out) {
    const	uint32_t	warp=warp_size(8,N);
    _half2float<<<N/8/warp,warp>>>(inp.p,out.p);
}

__global__	void	_float2char(float	*inp, char	*out) {
    uint32_t	id=(blockIdx.x*blockDim.x+threadIdx.x)<<3;
    *(char8*)(out+id)=_f2c8(*(float8*)(inp+id));
}

template<uint32_t	N>
static	inline	void	float2char(tensor<float,N>	&inp,	tensor<char,N>	&out) {
    const	uint32_t	warp=warp_size(8,N);
    _float2char<<<N/8/warp,warp>>>(inp.p,out.p);
}

__global__	void	_char2float(char	*inp, float	*out) {
    uint32_t	id=(blockIdx.x*blockDim.x+threadIdx.x)<<3;
    *(float8*)(out+id)=_c2f8(*(char8*)(inp+id));
}

template<uint32_t	N>
static	inline	void	char2float(tensor<char,N>	&inp,	tensor<float,N>	&out) {
    const	uint32_t	warp=warp_size(8,N);
    _char2float<<<N/8/warp,warp>>>(inp.p,out.p);
}

#endif
